Skip to content

[experimental] rocSHMEM as an allocation provider - #550

Open
nirvedhmeshram wants to merge 16 commits into
ROCm:mainfrom
nirvedhmeshram:nmeshram/rocshmem-provider
Open

nirvedhmeshram wants to merge 16 commits into
ROCm:mainfrom
nirvedhmeshram:nmeshram/rocshmem-provider

Conversation

@nirvedhmeshram

@nirvedhmeshram nirvedhmeshram commented Sep 8, 2026

Copy link
Copy Markdown

Summary (human)

Uses rocSHMEM API to do allocations and form map suggested in #546 , no changes to iris kernels are needed.

Motivation

Issue #546 proposes an allocator-agnostic boundary so Iris device kernels can operate on tensors from providers Iris does not own. This is a second implementation of that shape, against rocSHMEM, to test whether it holds.

The main result is stronger than expected: Iris device code needs no changes at all. iris.store/load/copy already take heap_bases as a plain pointer argument, and __translate computes offset = ptr - bases[from]; bases[to] + offset. Any table satisfying peer_bases[local_rank] == local allocation base drives them. The test kernel calls unmodified iris.store on memory allocated entirely by rocSHMEM, with no Iris context and no Iris heap anywhere in the process.

So the entire integration surface is host-side, and what a provider owes Iris is exactly one int64 table.

Technical details

iris/experimental/rocshmem_provider.py adds:

The table is built from rocshmem_ptr(base, peer), which is OpenSHMEM's shmem_ptr: an address in our own address space for the peer's counterpart, or NULL when that peer is not reachable by direct load/store.

One context-wide table serves every allocation. rocshmem_ptr is a single linear translation of the whole symmetric heap — GDAHostContext::shmem_ptr computes ipc_bases[peer] + (p - ipc_bases[me]) — so the peer delta is constant for every heap address regardless of allocation. Any symmetric anchor yields a table valid for all of them, which also means the provider never needs rocSHMEM's heap base (not exposed publicly). This is verified rather than assumed: a table built from one allocation is used to translate pointers belonging to another, and the data lands correctly.

That property matters for Iris specifically, because iris.copy takes a single heap_bases and translates two pointers against it. A provider handing out genuinely per-allocation tables could not drive it.

Relationship to #549

Compatible, not dependent. This targets main and was validated at 6432c101 with #549 not applied — it needs nothing from that PR, and touches no file it touches. allocate_symmetric mirrors its signature so the two line up when it lands, and will follow whatever shape it settles on.

Test plan

tests/unittests/test_rocshmem_provider.py, in the repo's pytest-under-torchrun convention:

python tests/run_tests_distributed.py tests/unittests/test_rocshmem_provider.py --num_ranks 2

Skips when rocshmem4py is absent, when fewer than 2 ranks are present, or when peers are not directly addressable. tests/manual_rocshmem_provider.py covers the multi-node case the unit test skips.

CI does not install rocSHMEM, so these skip there: the Test unittests jobs report collected 4 items / 4 skipped at 1, 2, 4 and 8 ranks and pass. The skip is applied in a fixture rather than at module scope on purpose — a module-level importorskip collects zero items, and pytest then returns exit code 5 (NO_TESTS_COLLECTED), which run_tests_distributed.py propagates and torchrun reports as a child failure, failing the whole job.

Run against real hardware with rocSHMEM installed (2+ ranks, one node, rocSHMEM built with USE_IPC=ON) the four tests execute and pass.

Note iris/experimental/__init__.py does not import the provider module, so import iris does not require rocshmem4py.

Test results

On MI355X (gfx950, ROCm 7.14, rocSHMEM 3.7.0 GDA/IONIC+IPC):

Run Result
test_rocshmem_provider.py, 2 ranks under the repo launcher 4 passed on both ranks, no skips
manual, 2 ranks 1 node PASS — unmodified iris.store over rocSHMEM memory
manual, cross-allocation PASS — a table from one allocation translates another's pointers
manual, 4 ranks across 2 nodes PASS — non-addressable peers detected on all 4 ranks

Not covered

  • Inter-node peers are detected and refused, not driven. Intra-node (IPC) only.
  • Only iris.store is exercised. load, put, get and the atomics are single-translation and should behave identically, but are untested here.
  • No views, no mixing two providers in one kernel, no imported tensors — three cases [Feature]: Formalize allocator-agnostic symmetric tensor address translation #546 calls out that this does not reach.

@nirvedhmeshram
nirvedhmeshram force-pushed the nmeshram/rocshmem-provider branch from 5abb6e6 to a28ffe6 Compare September 8, 2026 18:21
@nirvedhmeshram
nirvedhmeshram marked this pull request as ready for review September 8, 2026 20:41
Copilot AI lite review requested due to automatic review settings September 8, 2026 20:41

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Pull request overview

Adds an experimental rocSHMEM-backed “allocation provider” so existing Iris Triton device kernels can operate on rocSHMEM-allocated symmetric memory by supplying a compatible peer_bases table / address map.

Changes:

  • Introduces iris/experimental/rocshmem_provider.py with RocshmemProvider, allocate_symmetric(*), and SymmetricAddressMap.
  • Adds distributed pytest coverage that drives unmodified iris.store over rocSHMEM memory (skipping cleanly when unavailable).
  • Adds a manual torchrun script to exercise intra-node (IPC) and multi-node indirect-peer detection.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 10 comments.

File Description
iris/experimental/rocshmem_provider.py Implements rocSHMEM allocations and produces Iris-compatible peer-base tables / richer address descriptor.
tests/unittests/test_rocshmem_provider.py Adds distributed unit tests validating translation tables and iris.store over rocSHMEM memory.
tests/manual_rocshmem_provider.py Adds a manual launcher script to validate IPC and indirect-peer reporting outside CI.

💡 Add a code-review agent skill for context-aware, tailored reviews. Learn more in the docs.

Comment thread iris/experimental/rocshmem_provider.py
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread tests/manual_rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py Outdated
mawad-amd added a commit that referenced this pull request Sep 10, 2026
allocate_symmetric() now returns a table whose entry r is the address of
this tensor on rank r, so peer_bases[cur_rank] == tensor.data_ptr(). It
previously returned the heap base table, which named the heap rather than
the allocation.

This matches what the rocSHMEM provider in #550 returns, so the two line
up on the same integers and not merely the same signature. Device-side
translation is unchanged either way -- it subtracts peer_bases[cur_rank]
and adds peer_bases[to], and any consistent anchor works.

Computed on device: heap_bases[cur_rank] stays a tensor rather than going
through .item(), so there is no device-to-host sync on the allocation path.

The test's heap-membership assertion becomes an equality against
data_ptr(), which is the invariant the design actually rests on; the old
one-sided bound was cleared by any heap pointer.
Comment thread iris/experimental/rocshmem_provider.py Outdated
Comment thread iris/experimental/rocshmem_provider.py
Comment thread iris/experimental/rocshmem_provider.py
Comment thread iris/experimental/rocshmem_provider.py
Comment thread tests/manual_rocshmem_provider.py Outdated
Comment thread tests/manual_rocshmem_provider.py Outdated
Comment thread .github/scripts/install_rocshmem.sh Outdated
Comment thread docker/Dockerfile Outdated
nirvedhmeshram and others added 13 commits September 15, 2026 16:01
Lets Iris device code operate on tensors allocated by rocSHMEM instead of from
Iris's own symmetric heap. No Iris device code changes are required: store, load
and copy take heap_bases as a plain pointer argument, so any table satisfying
peer_bases[local_rank] == local allocation base drives them.

iris/experimental/rocshmem_provider.py builds that table from
rocshmem_ptr(base, peer), which returns an address in this process's own space
for a peer's counterpart of a symmetric object, or NULL when that peer is not
reachable by direct load/store.

  allocate_symmetric(*size, dtype)      -> (tensor, peer_bases)
  allocate_symmetric_map(*size, dtype)  -> (tensor, SymmetricAddressMap)
  symmetric_address_map(tensor)         -> SymmetricAddressMap

The first matches Iris.allocate_symmetric's shape so the same kernels drive
either provider. The descriptor form adds local_rank, allocation_base,
allocation_bytes and a per-peer `direct` mask; callers check that mask before
launching, since a peer that is not directly addressable has a base of 0 and
would translate to a wild pointer rather than an error.

One table serves every allocation. rocSHMEM's peer mapping is a linear
translation of the whole symmetric heap, so any symmetric address anchors a
table valid for all allocations, and rocSHMEM's heap base -- which it does not
expose publicly -- is never needed. That also keeps iris.copy usable, since it
translates two pointers against a single heap_bases.

Scope is intra-node. Inter-node peers are reported as unreachable rather than
driven; they need a transport this module does not provide.

Tests:
  tests/unittests/test_rocshmem_provider.py  pytest under the repo launcher,
      skipping when rocshmem4py is absent, when fewer than 2 ranks are present,
      or when peers are not directly addressable
  tests/manual_rocshmem_provider.py          multi-node script, including the
      non-addressable-peer path via EXPECT_INDIRECT=1

The provider module is not imported by iris/experimental/__init__.py, so
`import iris` does not require rocshmem4py.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
allocate_symmetric() cached the table built from the first allocation and
returned it for every later one. Translation still worked, because the per-peer
offset is constant across the heap, but peer_bases[local_rank] was the first
allocation's base rather than the current tensor's -- contradicting the stated
invariant and breaking the assertion in tests/manual_rocshmem_provider.py, which
checks it on a second allocation.

Cache the per-peer offsets instead and build each allocation's table from its
own base. peer_bases[local_rank] is now that allocation's base for every
allocation, while the shared offsets keep a table from one allocation able to
translate another's pointers, which iris.copy relies on.

This also removes a duplicated rocshmem_ptr sweep: the first allocation
previously queried every peer twice, once to build the map and once to seed the
cache. It is now queried once per process.

peer_bases is created on tensor.device rather than a device captured when the
provider was constructed, so the table cannot end up on a different device than
the memory it describes.

test_table_is_context_wide asserted the two tables were equal, which no longer
holds and was the weaker property anyway. It is now test_peer_offsets_are_shared
and checks what actually matters: each table's local entry is its own
allocation's base, and the per-peer offsets agree. Unreachable peers are
excluded from that comparison, since their entry is 0 rather than base + offset.

Verified on 2 ranks: pytest 4 passed, manual test PASS including the
cross-allocation check.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Review follow-ups, all documentation:

- allocate_symmetric: say it is collective, describe the returned table's
  shape, dtype, device and invariant, and point at allocate_symmetric_map
  for the direct mask.
- symmetric_address_map: note it builds a fresh table per call, why that is
  once per allocation rather than per launch, and why it is deliberately not
  memoised by data_ptr.
- free: explain it cannot be automated. rocshmem_free is documented as
  collective and must be called by all PEs, so a __del__ or weakref finalizer
  would let ranks diverge on GC timing and hang.
- module: record that rocshmem4py is a standalone package that statically
  links rocSHMEM rather than linking it at run time, so its USE_IPC setting
  is fixed at its build time.
The provider tests have only ever skipped in CI, because rocshmem4py was not
installed. Add it, so they actually execute.

Installing it is a source build but a short one. There is no prebuilt
rocshmem4py wheel on any index -- not PyPI, not the ROCm nightly indexes -- and
the bindings do not build rocSHMEM themselves (find_package(rocshmem 3.5.0
CONFIG REQUIRED), no FetchContent), so rocSHMEM is built first and the bindings
are pip-installed against it with CMAKE_PREFIX_PATH.

Scope is IPC only, which is what the provider uses and all a single-node runner
can exercise. Upstream already defaults USE_IPC=ON and USE_GDA=OFF, so no
conduit flags are passed, which keeps MPI and the RDMA provider libraries out of
it. Built for gfx942 to match the MI325X runners.

One installer shared by both image definitions rather than two copies. That
required widening the Docker build context from docker/ to the repo root so the
Dockerfile can COPY it, hence the .dockerignore; the Apptainer def pulls the same
file in via %files.

Also adds iris/experimental/README.md covering install, verification and the
collective-call contract, as requested in review.
Both are pip installs from source, but the README documented the git+ form while
the CI script used a local path, with nothing explaining the difference.

One checkout serves the core build and the bindings, so they are guaranteed to be
the same revision. A git+ URL would have pip clone the monorepo again at whatever
develop is at by then, and find_package would not catch the skew because it only
compares versions while the bindings statically link the core.
The docstring claimed the same signature and return shape as
Iris.allocate_symmetric. Iris has no such method -- its allocation API is
zeros/ones/full/uniform/as_symmetric, with a context-wide table from
get_heap_bases(). The (tensor, peer_bases) shape comes from the provider
interface being proposed, not from an existing method.
Comment only. The CI bases are ROCm 7.2.1 (apptainer) and 7.1 (docker), both
older than the 7.14 artifacts that ship rocSHMEM's static library and headers,
so the source build is still required. Records where to cut it when a base image
bumps, and that rocshmem4py stays a source build either way until its TheRock
packaging lands.
The provider tests now run in CI rather than skip, and rocSHMEM aborts inside
init with no message: a bare SIGABRT through library_init -> IPCBackend ->
Backend::init, 40ms in, with nothing to say why.

Set ROCSHMEM_DEBUG_LEVEL=info in both images so it prints its config banner, the
env vars it actually saw, and the backend it selected. Also write the rank, world
size and visible GPU count straight to fd 2 before the call, which survives both
pytest's capture and the abort, so the crash says what it was attempting.

Diagnostic only; drop both once the tests pass.
Review feedback.

The rationale for the source build, and for ROCSHMEM_DEBUG_LEVEL, was spelled out
in install_rocshmem.sh and then again in both image definitions. Keep it in the
script and leave a pointer at each call site.

Mark the ROCm 7.14 note as a TODO so it shows up when someone greps for work to
drop, rather than reading as background.
The previous diagnostic round produced nothing: neither rocSHMEM's banner nor a
write to fd 2 reached the CI log. pytest captures at the fd level and has already
redirected fd 2 by the time a fixture runs, so both landed in a buffer that is
discarded when the process aborts. faulthandler's output survived only because it
dups the original fd 2 at interpreter startup, which is why "Fatal Python error"
came through and nothing else did.

Suspend capture around init via the capturemanager plugin so rocSHMEM's own
logging, and the rank/GPU-count line, reach the log.

Also fold two overlapping "Returns (tensor, peer_bases)" paragraphs in
allocate_symmetric's docstring into one; they came from two separate edits.

Diagnostic; drop with the debug level once the tests pass.
With capture suspended, rocSHMEM's banner finally reached the log and named the
problem outright:

  # Compiled Arch(s)         : gfx942
  # System Arch              : gfx950:sramecc+:xnack-
  # System Arch is supported : No

The runner label is linux-mi325-8gpu-ossci-rad, and MI325X is gfx942, but the
hardware reports gfx950. rocSHMEM's device code has to match the GPU it runs on;
built for the wrong arch, the loaded code object has none of its device globals
and init aborts in HIP:

  hip_global.cpp:70 : Cannot create GlobalVar Obj for symbol:
  _ZN8rocshmem14logd_constantsE

Build both arches rather than swapping one guess for another, since the label
cannot be trusted to track the pool.
They did their job: the banner named the arch mismatch that 578394b fixes.

Removes ROCSHMEM_DEBUG_LEVEL=info from both images, and the fd-2 write plus the
pytest capture suspension from the provider fixture. The debug level printed a
config banner on every run, and the capture suspension is scaffolding that should
not outlive the bug it was added for.

Kept as its own commit so `git revert` restores the whole apparatus in one step
if a rocSHMEM init failure ever needs diagnosing again. Note that without it such
a failure is mute: pytest captures at the fd level, so rocSHMEM's output is
buffered and discarded when the process aborts.
ROCm#549 landed it on main, returning the same (tensor, peer_bases) pair with
peer_bases[cur_rank] == data_ptr(). An earlier commit dropped this reference
because the method did not exist yet; matching it is the whole point of the
provider, so say so.
@nirvedhmeshram
nirvedhmeshram force-pushed the nmeshram/rocshmem-provider branch from 4601a0c to 233c840 Compare September 15, 2026 21:02
A transient network failure during the Triton clone produced an image that
Apptainer reported as built, container_build.sh cached by def-file checksum, and
every later job reused:

  Cloning into '/opt/triton'...
  fetch-pack: unexpected disconnect while reading sideband packet
  fatal: early EOF
  /bin/bash: line 22: cd: /opt/triton: No such file or directory
  ERROR: file:///opt does not appear to be a Python project

The %post block had no `set -e`, so the failed clone, the failed cd, and the
failed editable install all continued. With /opt/triton absent and the pinned
3.7.0 checkout never installed, `import triton` fell back to the older
pytorch-triton-rocm in site-packages, which does not understand Python 3.14's
__annotate__. Every Triton and Gluon test then failed far downstream with
"Unsupported function referenced: <function ...__annotate__>", pointing nowhere
near the cause.

Add `set -e`, and retry the clone three times since it is large and this is a
transient failure that will recur. The Dockerfile is unaffected: each RUN already
fails on error.
$HOME is shared across the runners, so every job in a run uses one cache
directory. Jobs that start together all find no image and all build concurrently
into the same path -- observed in run 35023261876, where jobs on iris-mi350x-0
and iris-mi350x-2 both reported "Image or checksum not found" and built, while
jobs on -1, -2 and -3 read a checksum one of them had written.

Two problems with that. `apptainer build --force` writes IMAGE_PATH in place, so
a builder truncates the image other jobs are currently executing. And every job
that loses the race still pays for a full redundant build.

Take an flock around the check-and-build, and re-check inside it so whoever waits
usually finds the image already there. Build to a private temp path and rename
into place: rename is atomic and leaves the inode alone, so a job already running
the old image is unaffected.

Verified with two concurrent builders against a stub: one build, the other waited
and reused it, no temp files left. On build failure: exit 1, no checksum written,
no image published, temp removed.

@mawad-amd mawad-amd left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good. Thanks Nirvedh!

Review feedback: it duplicated the unit tests closely enough not to earn its
keep, and nothing ran it automatically -- it lives outside the directories CI
collects, its name does not match the test_*.py glob, and the launcher is
--nnodes=1 regardless.

Note what goes with it: the inter-node case, where rocshmem_ptr returns NULL and
SymmetricAddressMap.direct reports the peer as unreachable, now has no coverage
at all. The mask is still needed for correctness on more than one node. Said so
in the test module docstring and the README rather than leaving it implicit.

Also fix the build log while touching the file: every job printed "building new
Apptainer image" before taking the lock, including the ones that then found the
image already built. Now it says what each job actually did.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants